
/* Template project which demonstrates the basics on how to setup a project 
* example application.
* Host code.
*/

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes, project
#include <cutil_inline.h>
#include <cudpp/cudpp.h>


// includes, kernels
#include <template_kernel.cu>

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest( int argc, char** argv);

extern "C"
void computeGold( float* reference, float* idata, const unsigned int len);

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv) 
{
    runTest( argc, argv);

    cutilExit(argc, argv);
}

////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////
void
runTest( int argc, char** argv) 
{
	// use command-line specified CUDA device, otherwise use device with highest Gflops/s
	if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") )
		cutilDeviceInit(argc, argv);
	else
		cudaSetDevice( cutGetMaxGflopsDeviceId() );

    unsigned int timer = 0;
    cutilCheckError( cutCreateTimer( &timer));
    cutilCheckError( cutStartTimer( timer));
    
////////////////////////////////////////////////////////////////
//////  Set up parameters
    unsigned int numElements = 512;
    unsigned int num_threads = 32;
    //unsigned int mem_size = sizeof( float) * num_threads;
    unsigned int mem_size = sizeof( float) * numElements;

    // allocate host memory
    //float* h_idata = (float*) malloc( mem_size);
    float* h_idata = (float*) malloc( mem_size);
    // initalize the memory
for (unsigned int i = 0; i < numElements; ++i) 
    {
        h_idata[i] = (float) (rand() & 0xf);
    }

	for(int i=0;i<numElements;i++)
	{
	printf("%f, ", h_idata[i]);
	}
    // allocate device memory
    float *d_ipdata, *d_ipindex, *d_iptag,
		*d_imindex, *d_imlimit, *d_imvar;

    cutilSafeCall( cudaMalloc( (void**) &d_ipdata, mem_size));
    cutilSafeCall( cudaMalloc( (void**) &d_ipindex, mem_size));
    cutilSafeCall( cudaMalloc( (void**) &d_iptag, mem_size));
    cutilSafeCall( cudaMalloc( (void**) &d_imindex, mem_size));
    cutilSafeCall( cudaMalloc( (void**) &d_imlimit, mem_size));
    cutilSafeCall( cudaMalloc( (void**) &d_imvar, mem_size));
    // copy host memory to device
    cutilSafeCall( cudaMemcpy( d_ipdata, h_idata, mem_size,
                                cudaMemcpyHostToDevice) );
    cutilSafeCall( cudaMemcpy( d_iptag, h_idata, mem_size,
                                cudaMemcpyHostToDevice) );
  
	// allocate mem for the result on host side
    float* h_odata = (float*) malloc( mem_size);


	// initialize bsr
    // setup execution parameters
    dim3  grid( 16, 16);  // needs to be optimized
    dim3  threads( 16,16);
	//init_index<<< grid, threads, numElements >>>(d_ipindex, d_imindex, d_iptag, d_imlimit, d_imvar, numElements);

  

		// print 
		

    // allocate device memory for result
    float* d_odata;
    cutilSafeCall( cudaMalloc( (void**) &d_odata, mem_size));
///////////////////////////////////////////////////////////////////
	// CUDPP config
	
    CUDPPConfiguration config;
    config.op = CUDPP_ADD;
    config.datatype = CUDPP_FLOAT;
    config.algorithm = CUDPP_SORT_RADIX;
    config.options = CUDPP_OPTION_FORWARD | CUDPP_OPTION_EXCLUSIVE;
    
    CUDPPHandle scanplan = 0;
    CUDPPResult result = cudppPlan(&scanplan, config, numElements, 1, 0);  

    if (CUDPP_SUCCESS != result)
    {
        printf("Error creating CUDPPPlan\n");
        exit(-1);
    }
	

	// Run the scan
    //cudppScan (scanplan, d_odata, d_idata, numElements);
	
	// Run the sort
    cudppSort (scanplan, d_odata, d_ipdata, numElements);
//	radixSortFloatKeys (d_ipindex, d_ipdata, scanplan, numElements, false, 32);
/////////////////////////////////////////////////////////////////
	///// kernel stuff

    // setup execution parameters
    //dim3  grid( 1, 1, 1);
    //dim3  threads( num_threads, 1, 1);

    // execute the kernel
    //Shuffle_k<<< grid, threads, mem_size >>>( int* indexes, float* oldValues, float* 
	//					  newValues, 1, int numValues);
	Shuffle<<< grid, threads, numElements*2 >>>(d_ipdata, d_iptag, d_imvar);
    // check if kernel execution generated and error
    cutilCheckMsg("Kernel execution failed");

////////////////////////////////////////////
	printf("\nv Check Me v\n");
		// copy back 
	cutilSafeCall( cudaMemcpy( h_odata, d_iptag, sizeof( float) * numElements,
                                cudaMemcpyDeviceToHost) );
	printf("\n "); 
	for(int i=0;i<numElements;i++)
	{
	printf("%f, ", h_odata[i]);
	}
printf("\n^ Check Me ^\n");


/////////////////////////////////////////////////

		printf("\nv Check Me v\n");
// check initialization
		// copy back 


	cutilSafeCall( cudaMemcpy( h_odata, d_imvar, sizeof( float) * numElements,
                                cudaMemcpyDeviceToHost) );
	printf("\n "); 
	for(int i=0;i<numElements;i++)
	{
	printf("%f, ", h_odata[i]);
	}
printf("\n^ Check Me ^\n");

    // copy result from device to host
 //   cutilSafeCall( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads,
   //                             cudaMemcpyDeviceToHost) );


    // copy result from device to host
    cutilSafeCall( cudaMemcpy( h_odata, d_odata, sizeof( float) * numElements,
                                cudaMemcpyDeviceToHost) );
	printf("\n "); 
	for(int i=0;i<numElements;i++)
	{
	printf("%f, ", h_odata[i]);
	}

//////////////////end stuff/////////////////////////////////
	
    cutilCheckError( cutStopTimer( timer));
    printf( "\nProcessing time: %f (ms)\n", cutGetTimerValue( timer));
    cutilCheckError( cutDeleteTimer( timer));

    // compute reference solution
    //float* reference = (float*) malloc( mem_size);
    //computeGold( reference, h_idata, num_threads);

    // check result
	bool good=true;
   for (unsigned int i=1;i<numElements;i++)
   {
		//good+=(h_odata[i-1]<h_odata[i]);
		//if(good) good=(h_odata[i-1]<=h_odata[i]);
		if (h_odata[i-1]>h_odata[i])printf("%d, %d \n", h_odata[i-1], h_odata[i]);
   }
	if (good)printf("Good!\n");
	else 
   printf("\nWell? %d \n", good);
    // cleanup memory
    free( h_idata);
    free( h_odata);
    //free( reference);
    cutilSafeCall(cudaFree(d_ipdata));
    cutilSafeCall(cudaFree(d_ipindex));
    cutilSafeCall(cudaFree(d_iptag));
    cutilSafeCall(cudaFree(d_imindex));
    cutilSafeCall(cudaFree(d_imlimit));
    cutilSafeCall(cudaFree(d_imvar));
    cutilSafeCall(cudaFree(d_odata));

    cudaThreadExit();
}
